/**********************************************************************
Copyright �2013 Advanced Micro Devices, Inc. All rights reserved.

Redistribution and use in source and binary forms, with or without modification, are permitted provided that the following conditions are met:

�   Redistributions of source code must retain the above copyright notice, this list of conditions and the following disclaimer.
�   Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following disclaimer in the documentation and/or
 other materials provided with the distribution.

THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
 WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY
 DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS
 OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING
 NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
********************************************************************/


#include "SineWave.hpp"

char kernel_sinewave_func[] = "kernel_sinewave";
char kernel_func[32] = "kernel_add_func";

/* specify "__attribute__((reqd_work_group_size(" ATOMIC_WG_SIZE_STR ", 1, 1)))"
 because of the BRN where the kernel enqueues extra instances. */
static char g_pszAtomicBinaryTemplate[] =
{
	"//__attribute__((reqd_work_group_size(%d, 1, 1)))\n"
	"kernel void GlobalAtomicTest(__global int* in, __global int* out)\n"
	"{\n"
	"\tint gid = get_global_id(0);\n"
	"\tout[gid] = atomic_%s(in, 2);\n"
	"}\n"
	"\n"
	"//__attribute__((reqd_work_group_size(%d, 1, 1)))\n"
	"kernel void LocalAtomicTest(__global int* in, __global int* out)\n"
	"{\n"
	"\tlocal int loc;\n"
	"\tint gid = get_global_id(0);\n"
	"\tif (get_local_id(0) == 0) loc = 0;\n"
	"\tbarrier(CLK_LOCAL_MEM_FENCE);\n"
	"\tout[gid] = atomic_%s(&loc, 2);\n"
	"}\n"
	"\n"
};
int SineWave::setupSineWave()
{
    cl_uint inputSizeBytes;

    // allocate and init memory used by host
	std::cout << " func:" << __func__ << " line:" << __LINE__ << " width:" << width << " height:" << height << std::endl;
    inputSizeBytes = width * height * sizeof(cl_uint);
    input  = (cl_uint *) malloc(inputSizeBytes);

    CHECK_ALLOCATION(input, "Failed to allocate host memory. (input)");
    inputSizeBytes = width * height * sizeof(cl_float4);
    output = (cl_float  *) malloc(inputSizeBytes);
    CHECK_ALLOCATION(output, "Failed to allocate host memory. (output)");

	std::cout << " func:" << __func__ << " line:" << __LINE__ << " maskWidth:" << maskWidth << " maskHeight:" << maskHeight << std::endl;
    cl_uint maskSizeBytes = maskWidth * maskHeight * sizeof(cl_float);
    mask = (cl_float  *) malloc(maskSizeBytes);
    CHECK_ALLOCATION(mask, "Failed to allocate host memory (mask)");
#if 1
    for(cl_uint i = 0; i < width*height; i++)
    {
        input[i] = width*height - i;
    }
    for(cl_uint i = 0; i < maskWidth*maskHeight; i++)
    {
        mask[i] = 1;
    }
#else
    // random initialisation of input
    fillRandom<cl_uint >(input, width, height, 0, 255);

    // Fill a blurr filter or some other filter of your choice
    for(cl_uint i = 0; i < maskWidth*maskHeight; i++)
    {
        mask[i] = 0;
    }

    cl_float val = 1.0f / (maskWidth * 2.0f - 1.0f);

    for(cl_uint i = 0; i < maskWidth; i++)
    {
        cl_uint y = maskHeight / 2;
        mask[y * maskWidth + i] = val;
    }

    for(cl_uint i = 0; i < maskHeight; i++)
    {
        cl_uint x = maskWidth / 2;
        mask[i * maskWidth + x] = val;
    }
#endif

    // Unless quiet mode has been enabled, print the INPUT array.
    if(!sampleArgs->quiet)
    {
        #if 0
        printArray<cl_uint >(
            "Original Input",
            input,
            width,
            height);
        #endif
    }
    
    return SDK_SUCCESS;
}


int
SineWave::genBinaryImage()
{
    bifData binaryData;
	std::cout << " func:" << __func__ << " line:" << __LINE__ << std::endl;
    binaryData.kernelName = std::string("SineWave_Kernels.cl");
    binaryData.flagsStr = std::string("");
    if(sampleArgs->isComplierFlagsSpecified())
    {
        binaryData.flagsFileName = std::string(sampleArgs->flags.c_str());
    }

    binaryData.binaryName = std::string(sampleArgs->dumpBinary.c_str());
    int status = generateBinaryImage(binaryData);
    return status;
}

/**
 * display devices
 * displays the devices in a platform
 * @param platform cl_platform_id
 * @param deviceType deviceType
 * @return 0 if success else nonzero
 */
static int displayDevices_WORK_GROUP_SIZE(cl_platform_id platform, cl_device_type deviceType)
{
    cl_int status;
    // Get platform name
    char platformVendor[1024];
    status = clGetPlatformInfo(platform, CL_PLATFORM_VENDOR, sizeof(platformVendor),
                               platformVendor, NULL);
    CHECK_OPENCL_ERROR(status, "clGetPlatformInfo failed");
    std::cout << "\nSelected Platform Vendor : " << platformVendor << std::endl;
    // Get number of devices available
    cl_uint deviceCount = 0;
    status = clGetDeviceIDs(platform, deviceType, 0, NULL, &deviceCount);
    CHECK_OPENCL_ERROR(status, "clGetDeviceIDs failed");
    cl_device_id* deviceIds = (cl_device_id*)malloc(sizeof(cl_device_id) *
                              deviceCount);
    CHECK_ALLOCATION(deviceIds, "Failed to allocate memory(deviceIds)");
    // Get device ids
    status = clGetDeviceIDs(platform, deviceType, deviceCount, deviceIds, NULL);
    CHECK_OPENCL_ERROR(status, "clGetDeviceIDs failed");
    // Print device index and device names
    for(cl_uint i = 0; i < deviceCount; ++i)
    {
        char deviceName[1024];
        status = clGetDeviceInfo(deviceIds[i], CL_DEVICE_NAME, sizeof(deviceName),
                                 deviceName, NULL);
        CHECK_OPENCL_ERROR(status, "clGetDeviceInfo failed");
        size_t wg_size = 0;
        status = clGetDeviceInfo(deviceIds[i], CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(size_t),
                                 &wg_size, NULL);
        CHECK_OPENCL_ERROR(status, "clGetDeviceInfo failed");

        std::cout << "Device " << i << " : " << deviceName
                  <<" Device ID is "<<deviceIds[i] << "  wg_size=" << wg_size << std::endl;
    }
    free(deviceIds);
    return SDK_SUCCESS;
}

int
SineWave::setupCL(void)
{
    cl_int status = 0;
    cl_device_type dType;
	std::cout << " func:" << __func__ << " line:" << __LINE__ << std::endl;

    if(sampleArgs->deviceType.compare("cpu") == 0)
    {
        dType = CL_DEVICE_TYPE_CPU;
    }
    else //deviceType = "gpu"
    {
        dType = CL_DEVICE_TYPE_GPU;
        if(sampleArgs->isThereGPU() == false)
        {
            std::cout << "GPU not found. Falling back to CPU device" << std::endl;
            dType = CL_DEVICE_TYPE_CPU;
        }
    }

    /*
     * Have a look at the available platforms and pick either
     * the AMD one if available or a reasonable default.
     */
    cl_platform_id platform = NULL;
    int retValue = getPlatform(platform, sampleArgs->platformId,
                               sampleArgs->isPlatformEnabled());
    CHECK_ERROR(retValue, SDK_SUCCESS, "getPlatform() failed");

    // Display available devices.
    retValue = displayDevices(platform, dType);
    CHECK_ERROR(retValue, SDK_SUCCESS, "displayDevices() failed");
    displayDevices_WORK_GROUP_SIZE(platform, dType);
    CHECK_ERROR(retValue, SDK_SUCCESS, "displayDevices_WORK_GROUP_SIZE() failed");

    // If we could find our platform, use it. Otherwise use just available platform.
    cl_context_properties cps[3] =
    {
        CL_CONTEXT_PLATFORM,
        (cl_context_properties)platform,
        0
    };

    context = clCreateContextFromType(
                  cps,
                  dType,
                  NULL,
                  NULL,
                  &status);

    CHECK_OPENCL_ERROR( status, "clCreateContextFromType failed.");

    // getting device on which to run the sample
    status = getDevices(context, &devices, sampleArgs->deviceId,
                        sampleArgs->isDeviceIdEnabled());
    CHECK_ERROR(status, SDK_SUCCESS, "getDevices() failed");

    {
        // The block is to move the declaration of prop closer to its use
        cl_command_queue_properties prop = 0;
        commandQueue = clCreateCommandQueueWithProperties  (
                           context,
                           devices[sampleArgs->deviceId],
                           &prop,
                           &status);
        CHECK_OPENCL_ERROR( status, "clCreateCommandQueueWithProperties  failed.");
    }

    inputBuffer = clCreateBuffer(
                      context,
                      CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
                      sizeof(cl_uint ) * width * height,
                      input,
                      &status);
    CHECK_OPENCL_ERROR(status, "clCreateBuffer failed. (inputBuffer)");
    std::cout << " func:" << __func__ << " line:" << __LINE__ << " outputBuffer sizeof(cl_uint ) * width * height=" << sizeof(cl_uint ) * width * height << std::endl;

    outputBuffer = clCreateBuffer(
                       context,
                       CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR,
                       sizeof(cl_float4 ) * width * height,
                       output,
                       &status);
    CHECK_OPENCL_ERROR( status,  "clCreateBuffer failed. (outputBuffer)");

    maskBuffer = clCreateBuffer(
                     context,
                     CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
                     sizeof(cl_float ) * maskWidth * maskHeight,
                     mask,
                     &status);
    CHECK_OPENCL_ERROR( status, "clCreateBuffer failed. (maskBuffer)");

    // create a CL program using the kernel source
    buildProgramData buildData;
    buildData.kernelName = std::string("SineWave_Kernels.cl");
    buildData.devices = devices;
    buildData.deviceId = sampleArgs->deviceId;
    buildData.flagsStr = std::string("");
    if(sampleArgs->isLoadBinaryEnabled())
    {
        buildData.binaryName = std::string(sampleArgs->loadBinary.c_str());
    }

    if(sampleArgs->isComplierFlagsSpecified())
    {
        buildData.flagsFileName = std::string(sampleArgs->flags.c_str());
    }

    retValue = buildOpenCLProgram(program, context, buildData);
    CHECK_ERROR(retValue, SDK_SUCCESS, "buildOpenCLProgram() failed");

    // get a kernel object handle for a kernel with the given name
    kernel = clCreateKernel(program, kernel_func, &status);
    CHECK_OPENCL_ERROR(status, "clCreateKernel failed.");

    return SDK_SUCCESS;
}

int SineWave::setWorkGroupSize()
{
    cl_int status = 0;
    status = kernelInfo.setKernelWorkGroupInfo(kernel,
             devices[sampleArgs->deviceId]);
    CHECK_ERROR(status, SDK_SUCCESS, " setKernelWorkGroupInfo() failed");

    globalThreads[0] = width*1;
    localThreads[0]  = width*1;
    globalThreads[1] = 1*height;
    localThreads[1]  = 1*height;
    std::cout << " func:" << __func__ << " line:" << __LINE__ << " globalThreads[0]=" << globalThreads[0] << " globalThreads[1]=" << globalThreads[1] << std::endl;
    std::cout << " func:" << __func__ << " line:" << __LINE__ << " localThreads[0]=" << localThreads[0] << " kernelInfo.kernelWorkGroupSize=" << kernelInfo.kernelWorkGroupSize << std::endl;

    if((cl_uint)(localThreads[0]) > kernelInfo.kernelWorkGroupSize)
    {
        if(!sampleArgs->quiet)
        {
            std::cout << "Out of Resources!" << std::endl;
            std::cout << "Group Size specified : "<< localThreads[0]
                      << std::endl;
            std::cout << "Max Group Size supported on the kernel : "
                      << kernelInfo.kernelWorkGroupSize <<std::endl;
            std::cout <<"Changing the group size to " << kernelInfo.kernelWorkGroupSize
                      << std::endl;
        }

        localThreads[0] = kernelInfo.kernelWorkGroupSize;
    }
    return SDK_SUCCESS;
}

int
SineWave::runCLKernels(void)
{
    cl_int   status;
    cl_event events[2];

    status = this->setWorkGroupSize();
    CHECK_ERROR(status, SDK_SUCCESS, "setWorkGroupSize() failed");

    // Set appropriate arguments to the kernel
    status = clSetKernelArg(
                 kernel,
                 0,
                 sizeof(cl_mem),
                 (void *)&outputBuffer);
    CHECK_OPENCL_ERROR( status, "clSetKernelArg failed. (outputBuffer)");

    cl_uint input_w = 8;
    status = clSetKernelArg(
                 kernel,
                 1,
                 sizeof(cl_uint),
                 (void *)&input_w);
    CHECK_OPENCL_ERROR( status, "clSetKernelArg failed. (inputBuffer)");

    cl_uint input_h = 8;
    status = clSetKernelArg(
                 kernel,
                 2,
                 sizeof(cl_uint),
                 (void *)&input_h);
    CHECK_OPENCL_ERROR( status, "clSetKernelArg failed. (maskBuffer)");

    cl_float input_time = 0;

    status = clSetKernelArg(
                 kernel,
                 3,
                 sizeof(cl_float),
                 (void *)&input_time);
    CHECK_OPENCL_ERROR( status, "clSetKernelArg failed. (inputDimensions)");

    std::cout << " func:" << __func__ << " line:" << __LINE__ << " globalThreads[0]=" << globalThreads[0] << std::endl;
    std::cout << " func:" << __func__ << " line:" << __LINE__ << " localThreads[0]=" << localThreads[0] << " kernelInfo.kernelWorkGroupSize=" << kernelInfo.kernelWorkGroupSize << std::endl;

    // Enqueue a kernel run call.
    status = clEnqueueNDRangeKernel(
                 commandQueue,
                 kernel,
                 2,
                 NULL,
                 globalThreads,
                 localThreads,
                 0,
                 NULL,
                 &events[0]);
    CHECK_OPENCL_ERROR( status, "clEnqueueNDRangeKernel failed.");

    status = clFlush(commandQueue);
    CHECK_OPENCL_ERROR(status,"clFlush() failed");

    status = waitForEventAndRelease(&events[0]);
    CHECK_ERROR(status, SDK_SUCCESS, "WaitForEventAndRelease(events[0]) Failed");

    // Enqueue readBuffer
    status = clEnqueueReadBuffer(
                 commandQueue,
                 outputBuffer,
                 CL_TRUE,
                 0,
                 width * height * sizeof(cl_float4),
                 output,
                 0,
                 NULL,
                 &events[1]);
    CHECK_OPENCL_ERROR( status, "clEnqueueReadBuffer failed.");

    status = clFlush(commandQueue);
    CHECK_OPENCL_ERROR(status,"clFlush() failed");

    status = waitForEventAndRelease(&events[1]);
    CHECK_ERROR(status, SDK_SUCCESS, "WaitForEventAndRelease(events[1]) Failed");
    return SDK_SUCCESS;
}

/**
 * Reference CPU implementation of Simple Convolution
 * for performance comparison
 */
void
SineWave::SineWaveCPUReference(cl_uint  *output,
        const cl_uint  *input,
        const cl_float *mask,
        const cl_uint  width,
        const cl_uint height,
        const cl_uint maskWidth,
        const cl_uint maskHeight)
{
    cl_uint vstep = (maskWidth  - 1) / 2;
    cl_uint hstep = (maskHeight - 1) / 2;

    // for each pixel in the input
    for(cl_uint x = 0; x < width; x++)
        for(cl_uint y = 0; y < height; y++)
        {
            /*
             * find the left, right, top and bottom indices such that
             * the indices do not go beyond image boundaires
             */
            cl_uint left    = (x           <  vstep) ? 0         : (x - vstep);
            cl_uint right   = ((x + vstep) >= width) ? width - 1 : (x + vstep);
            cl_uint top     = (y           <  hstep) ? 0         : (y - hstep);
            cl_uint bottom  = ((y + hstep) >= height)? height - 1: (y + hstep);

            /*
             * initializing wighted sum value
             */
            cl_float sumFX = 0;

            for(cl_uint i = left; i <= right; ++i)
                for(cl_uint j = top ; j <= bottom; ++j)
                {
                    /*
                     * performing wighted sum within the mask boundaries
                     */
                    cl_uint maskIndex = (j - (y - hstep)) * maskWidth  + (i - (x - vstep));
                    cl_uint index     = j                 * width      + i;

                    /*
                     * to round to the nearest integer
                     */
                    sumFX += ((float)input[index] * mask[maskIndex]);
                }
            sumFX += 0.5f;
            output[y*width + x] = cl_uint(sumFX);
        }
}

int SineWave::initialize()
{
    // Call base class Initialize to get default configuration
    if  (sampleArgs->initialize() != SDK_SUCCESS)
    {
        return SDK_FAILURE;
    }

    // Now add customized options
    Option* width_option = new Option;
    CHECK_ALLOCATION(width_option, "Memory allocation error.\n");
	std::cout << " func:" << __func__ << " line:" << __LINE__ << " width:" << width << " height:" << height << std::endl;

    width_option->_sVersion = "x";
    width_option->_lVersion = "width";
    width_option->_description = "Width of the input matrix";
    width_option->_type = CA_ARG_INT;
    width_option->_value = &width;
	std::cout << " func:" << __func__ << " line:" << __LINE__ << " width:" << width << " height:" << height << std::endl;

    sampleArgs->AddOption(width_option);
    delete width_option;

    Option* height_option = new Option;
    CHECK_ALLOCATION(height_option, "Memory allocation error.\n");

    height_option->_sVersion = "y";
    height_option->_lVersion = "height";
    height_option->_description = "Height of the input matrix";
    height_option->_type = CA_ARG_INT;
    height_option->_value = &height;
	std::cout << " func:" << __func__ << " line:" << __LINE__ << " width:" << width << " height:" << height << std::endl;

    sampleArgs->AddOption(height_option);
    delete height_option;

    Option* mask_width = new Option;
    CHECK_ALLOCATION(mask_width, "Memory allocation error.\n");

    maskWidth = 3;
    mask_width->_sVersion = "m";
    mask_width->_lVersion = "masksize";
    mask_width->_description = "Width of the mask matrix";
    mask_width->_type = CA_ARG_INT;
    mask_width->_value = &maskWidth;
    sampleArgs->AddOption(mask_width);
    delete mask_width;

    Option* num_iterations = new Option;
    CHECK_ALLOCATION(num_iterations, "Memory allocation error.\n");

    num_iterations->_sVersion = "i";
    num_iterations->_lVersion = "iterations";
    num_iterations->_description = "Number of iterations for kernel execution";
    num_iterations->_type = CA_ARG_INT;
    num_iterations->_value = &iterations;

    sampleArgs->AddOption(num_iterations);
    delete num_iterations;
#define DEBUG_MODE
#ifdef DEBUG_MODE
    Option* math_en = new Option;
    CHECK_ALLOCATION(math_en, "set sinewave mode error.\n");
    sinewaveMode = 0;
    math_en->_sVersion = "M";
    math_en->_lVersion = "Mode";
    math_en->_description = "set sinewave mode";
    math_en->_type = CA_ARG_INT;
    math_en->_value = &sinewaveMode;
    sampleArgs->AddOption(math_en);
    delete math_en;
#endif
    std::cout << " ./SineWave.elf -x 1024 -y 1" << std::endl;
    std::cout << " ./SineWave.elf -x 1024 -y 1 -M 1  //call function PcmSineWWave" << std::endl;

    return SDK_SUCCESS;
}

int SineWave::setup()
{
	std::cout << " func:" << __func__ << " line:" << __LINE__ << " width:" << width << " height:" << height << std::endl;

    if(!isPowerOf2(width))
    {
        width = roundToPowerOf2(width);
    }
    if(!isPowerOf2(height))
    {
        height = roundToPowerOf2(height);
    }
	std::cout << " func:" << __func__ << " line:" << __LINE__ << " width:" << width << " height:" << height << std::endl;

    maskHeight = maskWidth;

    if(!(maskWidth%2))
    {
        maskWidth++;
    }
    if(!(maskHeight%2))
    {
        maskHeight++;
    }
	std::cout << " func:" << __func__ << " line:" << __LINE__ << " maskWidth:" << maskWidth << " maskHeight:" << maskHeight << std::endl;

    if (setupSineWave() != SDK_SUCCESS)
    {
        return SDK_FAILURE;
    }

    int timer = sampleTimer->createTimer();
    sampleTimer->resetTimer(timer);
    sampleTimer->startTimer(timer);

    if (setupCL() != SDK_SUCCESS)
    {
        return SDK_FAILURE;
    }

    sampleTimer->stopTimer(timer);

    setupTime = (cl_double)sampleTimer->readTimer(timer);

    return SDK_SUCCESS;
}


int SineWave::run()
{
    int timer = sampleTimer->createTimer();
    sampleTimer->resetTimer(timer);
    sampleTimer->startTimer(timer);

    std::cout << "Executing kernel for " << iterations <<
              " iterations" << std::endl;
    std::cout << "-------------------------------------------" << std::endl;

    for(int i = 0; i < iterations; i++)
    {
        // Arguments are set and execution call is enqueued on command buffer
        if (runCLKernels() != SDK_SUCCESS)
        {
            return SDK_FAILURE;
        }
    }

    sampleTimer->stopTimer(timer);
    totalKernelTime = (double)(sampleTimer->readTimer(timer)) / iterations;
    std::cout << " func:" << __func__ << " line:" << __LINE__ << std::endl;
#if 0
    if(!sampleArgs->quiet)
    {
        //printArray<cl_uint >("Output", output, width, height);
        {
            //std::cout << " func:" << __func__ << " line:" << __LINE__ << " output=" << std::endl;
            for(int i = 0; i < height; i++)
            {
                for(int j = 0; j < width; j++)
                {
                    //std::cout<< std::setw(4) <<output[i * width * 4 + j * 4 + 0]<<", ";
                    //std::cout<< std::setw(4) <<output[i * width * 4 + j * 4 + 1]<<", ";
                // std::cout<< std::setw(4) <<output[i * width * 4 + j * 4 + 2]<<", ";
                    //std::cout<< std::setw(4) <<output[i * width * 4 + j * 4 + 3]<<", ";
                }
                std::cout<<"\n";
            }
            std::cout<<"\n";
        }
    }
    #endif
    std::cout << " func:" << __func__ << " line:" << __LINE__ << std::endl;

    return SDK_SUCCESS;
}
//#define PCM_SINE_WAVE
float animate = 0.0f;
#define PI 3.1415926535

#define PCM_N 160000

#define N 1024 

#define M 10
float cl_dataR[N];
float cpu_dataR[N];
int
SineWave::PcmSineWWave(void)
{
    int f1 = 100;//   	# 信号的特征频率1
	int f2 = 300;//   # 信号的特征频率2
	int f3 = 500;//   # 信号的特征频率3
    FILE *fp;
	float SinIn[PCM_N];
	float dataR[PCM_N],dataI[PCM_N];
	float dataA[PCM_N];

    float cl_dataR[PCM_N];
    float cpu_dataR[PCM_N];
	float db_v = 0;
	float db = pow(10, -3/20);
//sample/every second
	float  framerate = 32000;
//channel_num
	float channel_num = 1;
//bytes needed every sample
	float sample_width = 2;
	float bits_width = sample_width*8;
// seconds, long of data
	float duration = 5;
// frequeny of sinewave
	float sinewav_frequency_l = 1000;
// max value of samples, depends on bits_width
	float  max_val = pow(2, (bits_width-1)) - 1;
    printf("[%s %d]max_val : =%f\n", __func__, __LINE__, max_val);
    printf("[%s %d]db : =%f\n", __func__, __LINE__, db);
    std::cout << " func:" << __func__ << " line:" << __LINE__ << std::endl;

//#volume = 32767*db #0x7FFF=32767
    float volume = max_val*db ;//#2**(bits_width-1) - 1;
//#多个声道生成波形数据

    float num=duration*framerate;

    printf("[%s %d]volume : =%f\n", __func__, __LINE__, volume);
    printf("[%s %d]num : =%f\n", __func__, __LINE__, num);

	for(int i = 0; i < PCM_N; i++)
	{
		float j = i;
        SinIn[i] = 0 + volume * sin(2*PI*sinewav_frequency_l*(j*1/PCM_N)) + 0 * sin(2*PI*f2* (j*1/PCM_N)) + 0 * sin(2*PI*f3* (j*1/PCM_N));

        if(int(max_val) <= int(SinIn[i]))
        {
            max_val = SinIn[i];
        }
		cpu_dataR[i] = SinIn[i];
		dataI[i] = 0;
		dataA[i] = 0;
	}
    printf("[%s %d]pcm sine wave max_val : =%f\n", __func__, __LINE__, max_val);
	if((fp=fopen("pcm_input.txt","wb"))==NULL)
	{
 		printf("fopen input.txt failed\n");
 		exit(0);
	}
	for(int i = 0; i < PCM_N ; i++)
		fprintf(fp,"%f,",cpu_dataR[i]);
	fclose(fp);

    return 0;
}
int
SineWave::SimpleGLCPUReference(void)
{
    int f1 = 100;//   	# 信号的特征频率1
	int f2 = 300;//   # 信号的特征频率2
	int f3 = 500;//   # 信号的特征频率3
    FILE *fp;
	float SinIn[N];
	float dataR[N],dataI[N];
	float dataA[N];

	float db_v = 0;
	float db = pow(10, -3/20);
//sample/every second
	float  framerate = 32000;
//channel_num
	float channel_num = 1;
//bytes needed every sample
	float sample_width = 2;
	float bits_width = sample_width*8;
// seconds, long of data
	float duration = 5;
// frequeny of sinewave
	float sinewav_frequency_l = 1000;
// max value of samples, depends on bits_width
	float  max_val = pow(2, (bits_width-1)) - 1;
    printf("max_val : =%f\n", max_val);
    printf("db : =%f\n", db);

	for(int i = 0; i < N; i++)
	{
		float j = i;
		SinIn[i] = 10 + 7 * sin(2*PI*f1*(j*1/N)) + 5 * sin(2*PI*f2* (j*1/N)) + 3 * sin(2*PI*f3* (j*1/N));

        if(int(max_val) <= int(SinIn[i]))
        {
            max_val = SinIn[i];
        }
		cpu_dataR[i] = SinIn[i];
		dataI[i] = 0;
		dataA[i] = 0;
	}
    printf("common sinewave max_val : =%f\n", max_val);
	if((fp=fopen("cpu_input.txt","wb"))==NULL)
	{
 		printf("fopen cpu_input.txt failed\n");
 		exit(0);
	}
	for(int i = 0; i < N ; i++)
		fprintf(fp,"%f,",cpu_dataR[i]);
	fclose(fp);

    for(unsigned int i = 0; i < height; ++i)
    {
        for(unsigned int j = 0; j < width; ++j)
        {
            unsigned int x = j;
            unsigned int y = i;

            // calculate uv coordinates
            float u = x / (float)width;
            float v = y / (float)height;
            u = u * 2.0f - 1.0f;
            v = v * 2.0f - 1.0f;

            // calculate simple sine wave pattern
            float freq = 4.0f;
            float w = sin(u * freq + animate) * cos(v * freq + animate) * 0.5f;
            float h = j;
            u = 
            // write output vertex
            verificationOutput[i * width * 4 + j * 4 + 0] = u;
            //verificationOutput[i * width * 4 + j * 4 + 1] = w;
            //verificationOutput[i * width * 4 + j * 4 + 2] = v;
            //verificationOutput[i * width * 4 + j * 4 + 3] = 1.0f;
        }
    }

    return 0;
}
int SineWave::verifyResults()
{
    FILE *fp;
    verificationOutput = (cl_float *) malloc(width * height * sizeof(cl_float4 ));
    CHECK_ALLOCATION(verificationOutput,                         "Failed to allocate host memory. (verificationOutput)");
    memset(verificationOutput, 0, width * height * sizeof(cl_float4));
    SimpleGLCPUReference();
    std::cout << " func:" << __func__ << " line:" << __LINE__ << std::endl;
    //PcmSineWWave();
    // compare the results and see if they match   
    {
        int m = 0;
        float degree = 0.01f;
        for(int i = 0; i < height; i++)
        {
            for(int j = 0; j < width; j++)
            {
                cl_dataR[m++] = output[i * width * 4 + j * 4 + 0];
                for(int k = 0; k < width; k++)
                {
                    if(((verificationOutput[i * width * 4 + j * 4 + k] + degree) > output[i * width * 4 + j * 4 + k]) && (((verificationOutput[i * width * 4 + j * 4 + k] - degree) < output[i * width * 4 + j * 4 + k])))
                    {
                        //std::cout<< std::setw(4) <<verificationOutput[i * width * 4 + j * 4 + k]<<"\n";
                        //std::cout<< std::setw(4) <<output[i * width * 4 + j * 4 + k]<<"\n";
                        //std::cout<<" if (((verificationOutput[i * width * 4 + j * 4 + k] + degree) > output[i * width * 4 + j * 4 + k]) && (((verificationOutput[i * width * 4 + j * 4 + k] - degree) < output[i * width * 4 + j * 4 + k])))\n";
                    }
                    else
                    {
                       // std::cout<<" else (((verificationOutput[i * width * 4 + j * 4 + k] + degree) > output[i * width * 4 + j * 4 + k]) && (((verificationOutput[i * width * 4 + j * 4 + k] - degree) < output[i * width * 4 + j * 4 + k])))\n";
                    }
                }
            }
           // std::cout<<"\n";
        }
    }
	if((fp=fopen("cl_input.txt","wb"))==NULL)
	{
 		printf("fopen input.txt failed\n");
 		exit(0);
	}
	for(int i = 0; i < N ; i++)
		fprintf(fp,"%f,",cl_dataR[i]);
	//fwrite(&dataR,sizeof(dataR),1,fp);
	fclose(fp);
    printf("\n");
	for(int i = 0; i < N ; i++)
    {
        //printf("cpu_dataR[%d]=%f cpu_dataR[%d]=%f\n", i, cpu_dataR[i], i, cpu_dataR[i]);
    }
    printf("\n");
    return SDK_SUCCESS;
}

void SineWave::printStats()
{
    if(sampleArgs->timing)
    {
        std::string strArray[5] = {"Width", "Height", "mask Size", "Time(sec)", "KernelTime(sec)"};
        std::string stats[5];

        sampleTimer->totalTime = setupTime + totalKernelTime;

        stats[0]  = toString(width    , std::dec);
        stats[1]  = toString(height   , std::dec);
        stats[2]  = toString(maskWidth, std::dec);
        stats[3]  = toString(sampleTimer->totalTime, std::dec);
        stats[4]  = toString(totalKernelTime, std::dec);

        printStatistics(strArray, stats, 5);
    }
}

int SineWave::cleanup()
{
    // Releases OpenCL resources (Context, Memory etc.)
    cl_int status;

    status = clReleaseKernel(kernel);
    CHECK_OPENCL_ERROR(status, "clReleaseKernel failed.(kernel)");

    status = clReleaseProgram(program);
    CHECK_OPENCL_ERROR(status, "clReleaseProgram failed.(program)");

    status = clReleaseMemObject(inputBuffer);
    CHECK_OPENCL_ERROR(status, "clReleaseMemObject failed.(inputBuffer)");

    status = clReleaseMemObject(outputBuffer);
    CHECK_OPENCL_ERROR(status, "clReleaseMemObject failed.(outputBuffer)");

    status = clReleaseMemObject(maskBuffer);
    CHECK_OPENCL_ERROR(status, "clReleaseMemObject failed.(maskBuffer)");

    status = clReleaseCommandQueue(commandQueue);
    CHECK_OPENCL_ERROR(status, "clReleaseCommandQueue failed.(commandQueue)");

    status = clReleaseContext(context);
    CHECK_OPENCL_ERROR(status, "clReleaseContext failed.(context)");

    // release program resources (input memory etc.)
    FREE(input);
    FREE(output);
    FREE(mask);
    FREE(verificationOutput);
    FREE(devices);

    return SDK_SUCCESS;
}

int
main(int argc, char * argv[])
{
    SineWave clSineWave;
    std::cout << " func:" << __func__ << " line:" << __LINE__ << " sinewaveMode:" << clSineWave.sinewaveMode << std::endl;
    if (clSineWave.initialize() != SDK_SUCCESS)
    {
        return SDK_FAILURE;
    }

    if (clSineWave.sampleArgs->parseCommandLine(argc, argv) != SDK_SUCCESS)
    {
        return SDK_FAILURE;
    }

    std::cout << " func:" << __func__ << " line:" << __LINE__ << " sinewaveMode:" << clSineWave.sinewaveMode << std::endl;
    if(0 == clSineWave.sinewaveMode)
    {
        strcpy(kernel_func, kernel_sinewave_func);
        std::cout << " func:" << __func__ << " line:" << __LINE__ << " sinewaveMode:" << clSineWave.sinewaveMode << " kernel_func=" << kernel_func << std::endl;
    }
    if(1 == clSineWave.sinewaveMode)
    {
        strcpy(kernel_func, kernel_sinewave_func);
        std::cout << " func:" << __func__ << " line:" << __LINE__ << " sinewaveMode:" << clSineWave.sinewaveMode << " kernel_func=" << kernel_func << " then call PcmSineWWave()"<< std::endl;
        clSineWave.PcmSineWWave();
    }
    else
    {
        std::cout << " func:" << __func__ << " line:" << __LINE__ << " sinewaveMode:" << clSineWave.sinewaveMode << " kernel_func=" << kernel_func << std::endl;
    }
    std::cout << " func:" << __func__ << " line:" << __LINE__ << std::endl;

    if(clSineWave.sampleArgs->isDumpBinaryEnabled())
    {
        return clSineWave.genBinaryImage();
    }

    if (clSineWave.setup() != SDK_SUCCESS)
    {
        return SDK_FAILURE;
    }

    if (clSineWave.run() != SDK_SUCCESS)
    {
        return SDK_FAILURE;
    }

    if (clSineWave.verifyResults() != SDK_SUCCESS)
    {
        return SDK_FAILURE;
    }

    if (clSineWave.cleanup() != SDK_SUCCESS)
    {
        return SDK_FAILURE;
    }

    clSineWave.printStats();
    return SDK_SUCCESS;
}
